

/* Template project which demonstrates the basics on how to setup a project 
 * example application.
 * Device code.
 */

#ifndef _TEMPLATE_KERNEL_H_
#define _TEMPLATE_KERNEL_H_

#include <stdio.h>

#define SDATA( index)      cutilBankChecker(sdata, index)

/*
__global__ void Shuffle_k(uint* indexes, float* oldValues, float* 
newValues, uint numElements, uint numValues) 
{ 
  extern __shared__ float values[]; 
  uint idx = threadIdx.x + blockIdx.x * blockDim.x; 
  if (idx < numElements) // In case the last block is not full 
    { 
      for (uint v = 0; v < numValues; v++) 
        { 
          values[threadIdx.x] = oldValues[indexes[idx] + v * numElements]; 
          newValues [idx + v * numElements] = values[threadIdx.x]; 
        } 
    } 

} 
*/

////////////////////////////////////////////////////////////////////////////////
//! Simple test kernel for device functionality
//! @param g_idata  input data in global memory
//! @param g_odata  output data in global memory
////////////////////////////////////////////////////////////////////////////////
__global__ void Shuffle_k(int* indexes, float* oldValues, float* 
						  newValues, int numElements, int numValues)
{
extern __shared__ float values[]; 
  int idx = threadIdx.x + blockIdx.x * blockDim.x; 
  if (idx < numElements) // In case the last block is not full 
    { 
      for (int v = 0; v < numValues; v++) 
        { 
          values[threadIdx.x] = oldValues[indexes[idx] + v * numElements]; 
          newValues [idx + v * numElements] = values[threadIdx.x]; 
        } 
    } 
}

__global__ void Shuffle(float *indexes, float *oldvalues, float *newvalues)
{
	int idx = threadIdx.x + blockIdx.x * blockDim.x; 

	newvalues[idx]=oldvalues[(int)indexes[idx]];


}

__global__ void
testKernel( float* g_idata, float* g_odata) 
{
  // shared memory
  // the size is determined by the host application
  extern  __shared__  float sdata[];

  // access thread id
  const unsigned int tid = threadIdx.x;
  // access number of threads in this block
  const unsigned int num_threads = blockDim.x;

  // read in input data from global memory
  // use the bank checker macro to check for bank conflicts during host
  // emulation
  SDATA(tid) = g_idata[tid];
  __syncthreads();

  // perform some computations
  SDATA(tid) = (float) num_threads * SDATA( tid);
  __syncthreads();

  // write data to global memory
  g_odata[tid] = SDATA(tid);
}


#endif // #ifndef _TEMPLATE_KERNEL_H_


// initiailize BSR
__global__ void init_index(float *pindex, float *mindex,float *ptag,
						   float *mlimit, float *mvar, int N)
{
	int idx = blockIdx.x * blockDim.x + threadIdx.x;
	pindex[idx]=idx;
	mindex[idx]=idx+N;
	mvar[idx]=0;

	//ptag[idx]=0;
	mlimit[idx]=1;
}

